/* Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved.
 *
 * Redistribution and use in source and binary forms, with or without
 * modification, are permitted provided that the following conditions
 * are met:
 *  * Redistributions of source code must retain the above copyright
 *    notice, this list of conditions and the following disclaimer.
 *  * Redistributions in binary form must reproduce the above copyright
 *    notice, this list of conditions and the following disclaimer in the
 *    documentation and/or other materials provided with the distribution.
 *  * Neither the name of NVIDIA CORPORATION nor the names of its
 *    contributors may be used to endorse or promote products derived
 *    from this software without specific prior written permission.
 *
 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
 * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
 * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
 * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
 * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
 * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
 * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
 * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
 * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 */

/**
 * Vector addition: C = A + B.
 *
 * This sample is a very basic sample that implements element by element
 * vector addition. It is the same as the sample illustrating Chapter 2
 * of the programming guide with some additions like error checking.
 */

#include <stdio.h>
#include <sys/time.h>
// For the CUDA runtime routines (prefixed with "cuda_")
#include <cuda_runtime.h>

#include <helper_cuda.h>
/**
 * CUDA Kernel Device code
 *
 * Computes the vector addition of A and B into C. The 3 vectors have the same
 * number of elements numElements.
 */
__global__ void gpu_sc(const int *A, int *C, int w, int h) {
    int i = blockDim.x * blockIdx.x + threadIdx.x;

    if (i < w*h) {
        if(0 == ((i) % 4))
        {
            C[i/4] = A[i];
        }
    }
}
void operateMatrix(int *InputData, int *OutputData, int x, int y, int nComputeWidthNum, int nComputeHeightNum, int nBlockINum, int nBlockJNum)
{
    for (int i = 0; i < 0 + nComputeWidthNum; i++)
    {
        for (int j = 0; j < 0 + nComputeHeightNum; j++)
        {
            int row = (i + 0) * ((nBlockINum-x));
            int col = (j + 0);// * ((nBlockJNum-y));
            //printf("[%s +%d %s] (nBlockINum=%03d nBlockJNum=%03d)(%03d, %03d) (%03d, %03d)\n", __FILE__, __LINE__, __func__, x, y, i, j, row * nComputeWidthNum, col);
            printf("[%s +%d %s] (x=%03d y=%03d) (row=%03d col=%03d) (nBlockINum=%03d nBlockJNum=%03d) InputData[%d]=%03d\n", __FILE__, __LINE__, __func__, \
                   x, y, row, col, nBlockINum, nBlockJNum,
                   col + row * nComputeWidthNum, InputData[col + row * nComputeWidthNum]);
            OutputData[col + row * nComputeWidthNum] = InputData[col + row * nComputeWidthNum];
        }
        printf(">\n");
    }
    printf(">>\n");
}
void cpu_matrixsum(int *InputData, int InputXNum, int InputYNum, int *OutputData, int OutputXNum, int OutputYNum, int nComputeWidthNum, int nComputeHeightNum)
{
#if 1
#else
    int colIndex;
#if 1
    int i, j;

    for(i = 0; i < InputXNum; i=i+nComputeWidthNum) {
        for(j = 0; j < InputYNum; j=j+nComputeHeightNum) {
#if 0
            for (int k = i; k < i + nComputeWidthNum && k < InputXNum; k++) {
                for (int l = j; l < j + nComputeHeightNum && l < InputYNum; l++) {
                    //block.push_back(arr[k][l]);
                    int min = i + nComputeWidthNum;
                    if(min > InputXNum)
                    {
                        min = InputXNum;
                    }
                    OutputData[k * min + colIndex] = InputData[k * min + j];

                }
            }
#else
            int i, j, k;
            int nNum = 0;
            for(k = 0; k < nComputeWidthNum; k++) {
                int row = i * nComputeWidthNum + k;
                if(row >= InputXNum) break;
                for(int col = 0; col < nComputeHeightNum; col++) {
                    colIndex = j * nComputeHeightNum + col;
                    if(colIndex >= InputYNum) break;
                    OutputData[row * InputXNum + colIndex] = InputData[row * InputXNum + colIndex];
                    nNum ++;
                    printf("[%s +%d %s] nNum=%d row * InputXNum + colIndex=%d \n", __FILE__, __LINE__, __func__, nNum, row * InputXNum + colIndex);
                }
                //printf("\n");
            }
#endif
        }
    }
#else
    int nBlockINum = InputXNum / nComputeWidthNum;
    int nBlockJNum = InputYNum / nComputeHeightNum;
    printf("[%s +%d %s] InputXNum=%d, InputYNum=%d\n", __FILE__, __LINE__, __func__, InputXNum, InputYNum);
    printf("[%s +%d %s] OutputXNum=%d, OutputYNum=%d\n", __FILE__, __LINE__, __func__, OutputXNum, OutputYNum);
    printf("[%s +%d %s] nComputeWidthNum=%d, nComputeHeightNum=%d\n", __FILE__, __LINE__, __func__, nComputeWidthNum, nComputeHeightNum);
    printf("[%s +%d %s] nBlockINum=%d, nBlockHeightNum=%d\n", __FILE__, __LINE__, __func__, nBlockINum, nBlockJNum);
    for(int i = 0; i < nBlockINum; i++)
    {
        for(int j = 0; j < nBlockJNum; j++)
        {
            int x = i * 1;
            int y = j * 1;
            operateMatrix(InputData, OutputData, x, y, nComputeWidthNum, nComputeHeightNum, nBlockINum, nBlockJNum);
            //OutputData[row * InputXNum + col] = InputData[row * InputXNum + col];

        }
    }
#endif
#endif
}
void cpu_matrixdiv(int *InputData, int InputXNum, int InputYNum, int *OutputData, int OutputXNum, int OutputYNum, int nComputeWidthNum, int nComputeHeightNum)
{
    int nBlockINum = InputXNum / nComputeWidthNum;
    int nBlockJNum = InputYNum / nComputeHeightNum;
    printf("[%s +%d %s] InputXNum=%d, InputYNum=%d\n", __FILE__, __LINE__, __func__, InputXNum, InputYNum);
    printf("[%s +%d %s] OutputXNum=%d, OutputYNum=%d\n", __FILE__, __LINE__, __func__, OutputXNum, OutputYNum);
    printf("[%s +%d %s] nComputeWidthNum=%d, nComputeHeightNum=%d\n", __FILE__, __LINE__, __func__, nComputeWidthNum, nComputeHeightNum);
    printf("[%s +%d %s] nBlockINum=%d, nBlockJNum=%d\n", __FILE__, __LINE__, __func__, nBlockINum, nBlockJNum);
    for (int i = 0; i < InputXNum; i += nComputeWidthNum) {
        for (int j = 0; j < InputYNum; j += nComputeWidthNum) {
            for (int m = i; m < i + nComputeWidthNum && m < InputXNum; ++m) {

                for (int n = j; n < j + nComputeWidthNum && n < InputXNum; ++n) {

                    OutputData[m * InputXNum + n] = InputData[m * InputXNum + n];
                    //printf(" [%d][%d]=%04d  ", m, n, out[m][n]);

                }
                //printf("\n");
            }
            //printf("\n");

        }
    }

}
void cpu_sc_matrix(int *InputData, int InputXNum, int InputYNum, int *OutputData, int OutputXNum, int OutputYNum, int nComputeWidthNum, int nComputeHeightNum)
{
    for(int i = 0; i < InputYNum * InputXNum; i++)
    {
        if(0 == (i % 4))
        {
            *(OutputData + i/4) = *(InputData + i);
        }
        printf("[%s +%d %s] InputYNum=%d InputXNum=%d *(InputData + i=%d)=%d\n", __FILE__, __LINE__, __func__, InputYNum, InputXNum, i, *(InputData + i));
    }
}
void show_values(int input[], int nDataXNum, int nDataYNum)
{
    int nValue0Num = 0;
    //int nShowValuesNum = 64;
    int nDataNum = nDataXNum * nDataYNum;
   // nShowValuesNum = nDataXNum * nDataYNum;
    for (int i = 0; i < nDataNum; i++)
    {
        if(0 == input[i])
        {
            nValue0Num++;
        }
    }
    if(nDataNum == nValue0Num)
    {
        printf("[%s +%d %s] ----------------data value is 0!------------------", __FILE__, __LINE__, __func__);
    }
    else
    {
        for (int i = 0; i < nDataNum; i++)
        {
            //if(i < nShowValuesNum)
            {
                printf(" %04d ", input[i]);
                if(0 == ((i + 1) % nDataYNum))
                {
                    printf("\n");
                }
            }

        }
    }
    printf("\n");
}
/**
 * Host main routine
 */
int main() {
    // Error code to check return values for CUDA calls
    cudaError_t err = cudaSuccess;
    int bShowFlag = 0; // 0:NOT Show 1:Show
    int nInputWidthNum = 16;
    int nInputHeightNum = 16;
    int nOutputWidthNum = nInputWidthNum;
    int nOutputHeightNum = nInputHeightNum;
    int nComputeWidthNum = 8;
    int nComputeHeightNum = 4;
    nOutputWidthNum = nComputeWidthNum / 2;
    nOutputHeightNum = nComputeHeightNum / 2;
    printf("[%s +%d %s] nComputeWidthNum=%d nComputeHeightNum=%d\n", __FILE__, __LINE__, __func__, nComputeWidthNum, nComputeHeightNum);
    // Print the vector length to be used, and compute its nInputBytes
    int numElements = nInputWidthNum * nInputHeightNum;
    int nInputBytes = numElements * sizeof(int);
    int nComputeNum = nComputeWidthNum * nComputeHeightNum;
    int nComputeBytes = nComputeNum * sizeof(int);
    printf("[Vector addition of %d elements]\n", numElements);
    int nComputeCnt = numElements / nComputeNum;
    printf("[%s +%d %s] nComputeCnt=%d\n", __FILE__, __LINE__, __func__, nComputeCnt);
    int nOutputNum = nComputeNum / 4;
    int nOutputBytes = nComputeBytes / 4;
    printf("[%s +%d %s] nOutputNum=%d nOutputBytes=%d\n", __FILE__, __LINE__, __func__, nOutputNum, nOutputBytes);

    // Allocate the host input vector A
    int *pHostInput = (int *)malloc(nInputBytes);

    // Allocate the host output vector C
    int *pHostOutput = (int *)malloc(nOutputBytes);
    // Allocate the host output vector C
    int *pHostCompute = (int *)malloc(nComputeBytes);

    // Verify that allocations succeeded
    if (pHostInput == NULL || pHostOutput == NULL || pHostCompute == NULL) {
        fprintf(stderr, "Failed to allocate host vectors!\n");
        exit(EXIT_FAILURE);
    }
    bShowFlag = 1;
    // Initialize the host input vectors
    for (int i = 0; i < numElements; ++i) {
        pHostInput[i] = i;
    }
    for (int i = 0; i < nOutputNum; ++i) {
        pHostOutput[i] = 0;
    }
    for (int i = 0; i < nComputeNum; ++i) {
        pHostCompute[i] = 0;
    }

    // Allocate the device input vector A
    int *pDevInput = NULL;
    err = cudaMalloc((void **)&pDevInput, nInputBytes);

    if (err != cudaSuccess) {
        fprintf(stderr, "Failed to allocate device vector A (error code %s)!\n",
                cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    // Allocate the device output vector C
    int *pDevOutput = NULL;
    err = cudaMalloc((void **)&pDevOutput, nOutputBytes);
    if (err != cudaSuccess) {
        fprintf(stderr, "Failed to allocate device vector C (error code %s)!\n",
                cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }
    // Allocate the device output vector C
    int *pDevCompute = NULL;
    err = cudaMalloc((void **)&pDevCompute, nComputeBytes);
    if (err != cudaSuccess) {
        fprintf(stderr, "Failed to allocate device pDevCompute (error code %s)!\n",
                cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    if(1 == bShowFlag)
    {
        printf("[%s +%d %s] before cpu_matrixdiv show_values(pHostInput, dataXNum=%d, dataYNum=%d)\n", \
               __FILE__, __LINE__, __func__, nInputWidthNum, nInputHeightNum);
        show_values(pHostInput, nInputWidthNum, nInputHeightNum);
        printf("[%s +%d %s] before cpu_matrixdiv show_values(pHostOutput, dataXNum=%d, dataYNum=%d)\n", \
               __FILE__, __LINE__, __func__, nOutputWidthNum, nOutputHeightNum);
        show_values(pHostOutput, nOutputWidthNum, nOutputHeightNum);
        printf("[%s +%d %s] before cpu_matrixdiv show_values(pHostCompute, dataXNum=%d, dataYNum=%d)\n", \
               __FILE__, __LINE__, __func__, nComputeWidthNum, nComputeHeightNum);
        show_values(pHostCompute, nComputeWidthNum, nComputeHeightNum);
    }

    int nCompute_index = 0;
    for (int i = 0; i < nInputWidthNum; i += nComputeWidthNum) {
        for (int j = 0; j < nInputHeightNum; j += nComputeHeightNum) {
            int indexX = 0;
            int indexY = 0;
            for (int m = i; m < i + nComputeWidthNum && m < nInputWidthNum; ++m) {
                for (int n = j; n < j + nComputeHeightNum && n < nInputHeightNum; ++n) {
                    //printf(" [%d]=%04d (%d %d)=%04d", m * nInputWidthNum + n, pHostInput[m * nInputWidthNum + n], indexX, indexY, pHostCompute[indexX * nComputeWidthNum + indexY]);
                    //if(indexX < nComputeWidthNum && indexY < nComputeHeightNum)
                    {
                        pHostCompute[indexX * nComputeHeightNum + indexY] = pHostInput[m * nInputWidthNum + n];
                        //printf(" {[%d]=%04d (%d %d)=%04d j=%d n=%d indexX * nComputeHeightNum + indexY=%d} ", m * nInputWidthNum + n, \
                        //pHostInput[m * nInputWidthNum + n], indexX, indexY, pHostCompute[indexX * nComputeHeightNum + indexY], j, n, indexX * nComputeHeightNum + indexY);
                        indexY ++;
                    }
                    // else
                    {
                        //break;
                    }
                }
                //printf("\n");
                indexX ++;

                indexY = 0;
                //printf(">\n");
            }
            printf("<<\n");
            if(1 == bShowFlag)
            {
                printf("[%s +%d %s] before nCompute_index=%d show_values(pHostCompute, dataXNum=%d, dataYNum=%d)\n", __FILE__, __LINE__, __func__, \
                nCompute_index, nComputeWidthNum, nComputeHeightNum);
                show_values(pHostCompute, nComputeWidthNum, nComputeHeightNum);
            }                
            printf(">>\n");
            // Copy the host input vectors A and B in host memory to the device input
            // vectors in
            // device memory
            printf("Copy input data from the host memory to the CUDA device\n");
            err = cudaMemcpy(pDevCompute, pHostCompute, nComputeBytes, cudaMemcpyHostToDevice);

            if (err != cudaSuccess) {
                fprintf(stderr,
                        "Failed to copy vector A from host to device (error code %s)!\n",
                        cudaGetErrorString(err));
                exit(EXIT_FAILURE);
            }

            // Launch the Vector Add CUDA Kernel
            int threadsPerBlock = nComputeWidthNum * nComputeHeightNum;
            int blocksPerGrid = 1;//(numElements + threadsPerBlock - 1) / threadsPerBlock;
            printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid,
                threadsPerBlock);
            gpu_sc<<<blocksPerGrid, threadsPerBlock>>>(pDevCompute, pDevOutput, nComputeWidthNum, nComputeHeightNum);
            cudaDeviceSynchronize();
            err = cudaGetLastError();
            if (err != cudaSuccess) {
                fprintf(stderr, "Failed to launch gpu_sc kernel (error code %s)!\n",
                        cudaGetErrorString(err));
                exit(EXIT_FAILURE);
            }

            // Copy the device result vector in device memory to the host result vector
            // in host memory.
            printf("Copy output data from the CUDA device to the host memory\n");
            err = cudaMemcpy(pHostOutput, pDevOutput, nOutputBytes, cudaMemcpyDeviceToHost);

            if (err != cudaSuccess) {
                fprintf(stderr,
                        "Failed to copy vector C from device to host (error code %s)!\n",
                        cudaGetErrorString(err));
                exit(EXIT_FAILURE);
            }
            if(1 == bShowFlag)
            {
                printf("[%s +%d %s] after nCompute_index=%d show_values(pHostCompute, dataXNum=%d, dataYNum=%d)\n", \
                    __FILE__, __LINE__, __func__, nCompute_index, nComputeWidthNum, nComputeHeightNum);
                show_values(pHostCompute, nComputeWidthNum, nComputeHeightNum);
                printf("[%s +%d %s] after nCompute_index=%d show_values(pHostOutput, dataXNum=%d, dataYNum=%d)\n", \
                    __FILE__, __LINE__, __func__, nCompute_index, nOutputWidthNum, nOutputHeightNum);
                show_values(pHostOutput, nOutputWidthNum, nOutputHeightNum);
            }
            nCompute_index ++;

            //break;
        }
    }
    
    printf("Test PASSED\n");

    // Free device global memory
    err = cudaFree(pDevInput);

    if (err != cudaSuccess) {
        fprintf(stderr, "Failed to free device vector A (error code %s)!\n",
                cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    err = cudaFree(pDevOutput);

    if (err != cudaSuccess) {
        fprintf(stderr, "Failed to free device vector C (error code %s)!\n",
                cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    err = cudaFree(pDevCompute);

    if (err != cudaSuccess) {
        fprintf(stderr, "Failed to free device pDevCompute (error code %s)!\n",
                cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }

    // Free host memory
    free(pHostInput);
    free(pHostOutput);
    free(pHostCompute);

    printf("-----------------------------------------Done-----------------------------------------\n");
    return 0;
}
int cpu_writeint2file()
{
    int width = 640, height = 480;
    uint8_t yuv420pData[width*height*3];
    uint32_t total_bytes = width*height*3;
    uint32_t *data_ptr = NULL;
    int run_cnt = 0;
    int i = 0;
    char str_out[] = "./data_inters";
    FILE *outputFile = NULL;
    struct timeval begin, end;
    double elapsed = 0;
    printf("[%s +%d %s] width*height*3=%d total_bytes=%d str_out=%s\n", __FILE__, __LINE__, __func__, width*height*3, total_bytes, str_out);
    printf("[%s +%d %s] run_cnt=%d str_out=%s\n", __FILE__, __LINE__, __func__, run_cnt, str_out);

    outputFile = fopen(str_out, "wb");
    if(NULL == outputFile)
    {
        printf("[%s +%d %s] NULL == outputFile str_out=%s\n", __FILE__, __LINE__, __func__, str_out);
        return -1;
    }

    gettimeofday(&begin, NULL);
    data_ptr = (uint32_t *)(&yuv420pData[0]);
    printf("[%s +%d %s] total_bytes/4=%d\n", __FILE__, __LINE__, __func__, total_bytes/4);
    for(i = 0; i < total_bytes/4; i ++)
    {
      *(data_ptr + i) = (i + 0x5001);
      if (100 > i)
      {
        //printf("[%s +%d %s] *(data_ptr + %d)=0x%x\n", __FILE__, __LINE__, __func__, i, *(data_ptr + i));
      }
    }
    fwrite(yuv420pData, sizeof(yuv420pData), 1, outputFile);
    run_cnt ++;
    if(NULL != outputFile)
    {
        fclose(outputFile);
        outputFile = NULL;
    }
    gettimeofday(&end, 0);
    elapsed = (end.tv_sec - begin.tv_sec) +
                      ((end.tv_usec - begin.tv_usec) / 1000000.0);

    printf("[%s +%d %s] cpu_writeint2file total_bytes=%d elapsed=%fs\n", __FILE__, __LINE__, __func__, total_bytes, elapsed);

    return 0;
}
